Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

WIP Enable JIT compilation of CUDA code with root.exe #17317

Draft
wants to merge 3 commits into
base: master
Choose a base branch
from

Conversation

LukasBreitwieser
Copy link
Contributor

Working prototype of JIT compilation of CUDA code with bin/root.exe using the following test script:

.rawInput 1
__global__ void Kernel(int *out) { *out = 42; }
.rawInput 0
int *deviceOutput;
int hostOutput = 0;
cudaMalloc((void **) &deviceOutput, sizeof(int));
Kernel<<<1,1>>>(deviceOutput);
cudaGetLastError();
cudaDeviceSynchronize();
cudaMemcpy(&hostOutput, deviceOutput, sizeof(int), cudaMemcpyDeviceToHost);
std::cout << "Test " << (hostOutput == 42 ? "passed" : "FAILED") << std::endl;
.q

Remaining Issues

ROOT build without cpp modules

libcomplexDict.rootmap issue

Search for the root cause of the following error:

cling-ptx: Forward declarations from /home/lukas/sft/ws/root-project/build/Debug-no-modules/lib/libcomplexDict.rootmap:2:19: error: explicit specialization of undeclared template 
class 'complex'
template <> class complex<float>;
                  ^      ~~~~~~~
cling-ptx: Forward declarations from /home/lukas/sft/ws/root-project/build/Debug-no-modules/lib/libcomplexDict.rootmap:3:19: error: explicit specialization of non-template class '
complex'
template <> class complex<double>;

Quick fix: manually remove the following lines from libcomplexDict.rootmap

template <> class complex<float>;
template <> class complex<double>;
template <typename T> class _root_std_complex;

TClass error:

Processing hsimple.C...
cling-ptx: In file included from input_line_8:1:
In file included from /home/lukas/sft/ws/root-project/build/Debug-no-modules/include/TFile.h:28:
In file included from /home/lukas/sft/ws/root-project/build/Debug-no-modules/include/TDirectoryFile.h:25:
In file included from /home/lukas/sft/ws/root-project/build/Debug-no-modules/include/TDirectory.h:24:
In file included from /home/lukas/sft/ws/root-project/build/Debug-no-modules/include/TNamed.h:25:
In file included from /home/lukas/sft/ws/root-project/build/Debug-no-modules/include/TObject.h:18:
/home/lukas/sft/ws/root-project/build/Debug-no-modules/include/TStorage.h:85:4: error: incomplete type 'TClass' named in nested name specifier
   ClassDef(TStorage,0)  //Storage manager class

ROOT build with cpp modules

Fix wrong order of including headers:

cling: In file included from <built-in>:1:
In file included from /home/lukas/sft/ws/root-project/build/Debug/etc//cling/lib/clang/18/include/__clang_cuda_runtime_wrapper.h:157:
/home/lukas/sft/ws/root-project/build/Debug/etc//cling/lib/clang/18/include/__clang_cuda_device_functions.h:51:18: error: static declaration of '__cosf' follows non-static declarat
ion
__DEVICE__ float __cosf(float __a) { return __nv_fast_cosf(__a); }
                 ^
cling: /usr/include/x86_64-linux-gnu/bits/mathcalls.h:62:1: note: previous declaration is here
__MATHCALL_VEC (cos,, (_Mdouble_ __x));

TODOs:

  • Solve remaining issues described above
  • Add test script to roottest
  • Solve quick fixes in the code in a more fundamental way
  • Write more test cases

Copy link

Test Results

0 tests   0 ✅  0s ⏱️
0 suites  0 💤
0 files    0 ❌

Results for commit a60c2de.

@dpiparo dpiparo requested a review from guitargeek December 26, 2024 07:55
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

1 participant